__kernel void evalDepthMatch(
	__global float* observeDepthMap,
	__global float* renderDepthMap, 
	__global float* observeSkin,
	__global float* sum_1,
	__global float* sum_2,
	__global float* sum_3)
{
	int num = get_global_id(0);
	float sub = fabs(observeDepthMap[num]-renderDepthMap[num]);
	float rMatch;
	if(sub < 0.04f )
	{
		sum_1[num]=sub*100;
	}else{
		sum_1[num]=4;
	}

	if(sub < 0.01f && renderDepthMap[num] < 1.f){
		rMatch=1.0f;
	}else{
		rMatch=0.0f;
	}
	sum_2[num] = observeSkin[num] != 1.f || rMatch == 1.f;
	sum_3[num] = observeSkin[num] != 1.f && rMatch == 1.f;
}

__kernel void evalDepthMatch_Cheating(
	__global float* observeDepthMap,
	__global float* renderDepthMap, 
	__global float* observeSkin,
	__global float* resultDepth,
	__global float* resultSkin)
{
	int num = get_global_id(0);
	float sub = fabs(observeDepthMap[num]-renderDepthMap[num]);
	float sum_1=0;
	float sum_2=0;
	float sum_3=0;
	int rMatch=0;
	if(sub < 0.04f ){
		sum_1=sub;
	}else{
		sum_1=0.04f;
	}

	if(sub < 0.01f && renderDepthMap[num] < 1.f){
		rMatch=1;
	}
	//Cheating
	if(observeSkin[num] != 1.f || rMatch == 1) sum_2 = 0.0001220703125f;
	if(observeSkin[num] != 1.f && rMatch == 1) sum_3 = 0.00000001490116119384765625f;

	//Hacking
	resultDepth[num]=sum_1;
	resultSkin[num]=sum_2+sum_3;
}
//
//__kernel void evalDepthMatch_Cheating_Vector(
//	__global float4* observeDepthMap,
//	__global float4* renderDepthMap, 
//	__global float4* observeSkin,
//	__global float4* result,)
//{
//
//}

__kernel void reduction_scalar(
	__global float* data,
	__local float* partial_sums, 
	__global float* output) 
{
	int lid = get_local_id(0);
	int group_size = get_local_size(0);
	partial_sums[lid] = data[get_global_id(0)];
	barrier(CLK_LOCAL_MEM_FENCE);
	for(int i = group_size/2; i>0; i >>= 1) {
		if(lid < i) {
		partial_sums[lid] += partial_sums[lid + i];
		}
		barrier(CLK_LOCAL_MEM_FENCE);
	}
	if(lid == 0) {
	output[get_group_id(0)] = partial_sums[0];
	}
}

__kernel void reduction_vector(
	__global float4* data, 
    __local float4* partial_sums,
	__global float* output) {

   int lid = get_local_id(0);
   int group_size = get_local_size(0);

   partial_sums[lid] = data[get_global_id(0)];
   barrier(CLK_LOCAL_MEM_FENCE);

   for(int i = group_size/2; i>0; i >>= 1) {
      if(lid < i) {
         partial_sums[lid] += partial_sums[lid + i];
      }
      barrier(CLK_LOCAL_MEM_FENCE);
   }

   if(lid == 0) {
      output[get_group_id(0)] = dot(partial_sums[0], (float4)(1.0f));
   }
}

__kernel void reduction_complete(__global float4* data, 
      __local float4* partial_sums, __global float* sum) {

   int lid = get_local_id(0);
   int group_size = get_local_size(0);

   partial_sums[lid] = data[get_local_id(0)];
   barrier(CLK_LOCAL_MEM_FENCE);

   for(int i = group_size/2; i>0; i >>= 1) {
      if(lid < i) {
         partial_sums[lid] += partial_sums[lid + i];
      }
      barrier(CLK_LOCAL_MEM_FENCE);
   }

   if(lid == 0) {
      *sum = partial_sums[0].s0 + partial_sums[0].s1 +
             partial_sums[0].s2 + partial_sums[0].s3;
   }
}


